<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub::BlockScan&lt; T, BLOCK_THREADS, ALGORITHM &gt; Class Template Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li class="current"><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="annotated.html"><span>Class&#160;List</span></a></li>
      <li><a href="classes.html"><span>Class&#160;Index</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div id="nav-path" class="navpath">
  <ul>
<li class="navelem"><a class="el" href="namespacecub.html">cub</a></li><li class="navelem"><a class="el" href="classcub_1_1_block_scan.html">BlockScan</a></li>  </ul>
</div>
</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="classcub_1_1_block_scan-members.html">List of all members</a>  </div>
  <div class="headertitle">
<div class="title">cub::BlockScan&lt; T, BLOCK_THREADS, ALGORITHM &gt; Class Template Reference<div class="ingroups"><a class="el" href="group___block_module.html">Block-wide</a></div></div>  </div>
</div><!--header-->
<div class="contents">
<a name="details" id="details"></a><h2 class="groupheader">Detailed description</h2>
<div class="textblock"><h3>template&lt;
    typename T, 
    int BLOCK_THREADS, 
    BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt;<br/>
class cub::BlockScan&lt; T, BLOCK_THREADS, ALGORITHM &gt;</h3>

<p>The <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel prefix sum/scan of items partitioned across a CUDA thread block. </p>
<div class="image">
<img src="block_scan_logo.png" alt="block_scan_logo.png"/>
<div class="caption">
.</div></div>
 <dl class="section user"><dt>Overview</dt><dd>Given a list of input elements and a binary reduction operator, a <a href="http://en.wikipedia.org/wiki/Prefix_sum"><em>prefix scan</em></a> produces an output list where each element is computed to be the reduction of the elements occurring earlier in the input list. <em>Prefix sum</em> connotes a prefix scan with the addition operator. The term <em>inclusive</em> indicates that the <em>i</em><sup>th</sup> output reduction incorporates the <em>i</em><sup>th</sup> input. The term <em>exclusive</em> indicates the <em>i</em><sup>th</sup> input is not incorporated into the <em>i</em><sup>th</sup> output reduction.</dd></dl>
<dl class="section user"><dt></dt><dd>Optionally, <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> can be specialized by algorithm to accommodate different latency/throughput workload profiles:<ol type="1">
<li><b><a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c">cub::BLOCK_SCAN_RAKING</a></b>. An efficient "raking reduce-then-scan" prefix scan algorithm. <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab">More...</a></li>
<li><b><a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba7f51e58246eb53f1a97bd1bc8c0f400f">cub::BLOCK_SCAN_WARP_SCANS</a></b>. A quick "tiled warpscans" prefix scan algorithm. <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab">More...</a></li>
</ol>
</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">T</td><td>Data type being scanned </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">ALGORITHM</td><td><b>[optional]</b> <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab" title="BlockScanAlgorithm enumerates alternative algorithms for cub::BlockScan to compute a parallel prefix ...">cub::BlockScanAlgorithm</a> enumerator specifying the underlying algorithm to use (default: <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c">cub::BLOCK_SCAN_RAKING</a>)</td></tr>
  </table>
  </dd>
</dl>
<dl class="section user"><dt>A Simple Example</dt><dd>Every thread in the block uses the <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> class by first specializing the <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> type, then instantiating an instance with parameters for communication, and finally invoking collective member functions. </dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates an exclusive prefix sum of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], ..., [508,509,510,511] }</code>.</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Uses special instructions when applicable (e.g., warp <code>SHFL</code>)</li>
<li>Uses synchronization-free communication between warp lanes when applicable</li>
<li>Uses only one or two block-wide synchronization barriers (depending on algorithm selection)</li>
<li>Zero bank conflicts for most types</li>
<li>Computation is slightly more efficient (i.e., having lower instruction overhead) for:<ul>
<li>Prefix sum variants (<b><em>vs.</em></b> generic scan)</li>
<li>Exclusive variants (<b><em>vs.</em></b> inclusive)</li>
<li><code>BLOCK_THREADS</code> is a multiple of the architecture's warp size</li>
</ul>
</li>
<li>See <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab" title="BlockScanAlgorithm enumerates alternative algorithms for cub::BlockScan to compute a parallel prefix ...">cub::BlockScanAlgorithm</a> for performance details regarding algorithmic alternatives </li>
</ul>
</dd></dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00186">186</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>
</div><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html">TempStorage</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The operations exposed by <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the <code>__shared__</code> keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or <code>union</code>'d with other storage allocation types to facilitate memory reuse.  <a href="structcub_1_1_block_scan_1_1_temp_storage.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Methods</h2></td></tr>
<tr><td colspan="2"><div class="groupHeader">Collective constructors</div></td></tr>
<tr class="memitem:a982e1407d00b704c3046cd72c48acabb"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="a982e1407d00b704c3046cd72c48acabb"></a>
__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb">BlockScan</a> ()</td></tr>
<tr class="memdesc:a982e1407d00b704c3046cd72c48acabb"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <code>threadIdx.x</code>. <br/></td></tr>
<tr class="separator:a982e1407d00b704c3046cd72c48acabb"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ad0decc1ea510cd9c7df4cd380a26b1b1"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ad0decc1ea510cd9c7df4cd380a26b1b1">BlockScan</a> (<a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage)</td></tr>
<tr class="memdesc:ad0decc1ea510cd9c7df4cd380a26b1b1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>.  <a href="#ad0decc1ea510cd9c7df4cd380a26b1b1">More...</a><br/></td></tr>
<tr class="separator:ad0decc1ea510cd9c7df4cd380a26b1b1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ab364204badf83769f9140d799cc188a7"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ab364204badf83769f9140d799cc188a7">BlockScan</a> (int linear_tid)</td></tr>
<tr class="memdesc:ab364204badf83769f9140d799cc188a7"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#ab364204badf83769f9140d799cc188a7">More...</a><br/></td></tr>
<tr class="separator:ab364204badf83769f9140d799cc188a7"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a6264c3383db3384184cb13f5f282f73d"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a6264c3383db3384184cb13f5f282f73d">BlockScan</a> (<a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage, int linear_tid)</td></tr>
<tr class="memdesc:a6264c3383db3384184cb13f5f282f73d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#a6264c3383db3384184cb13f5f282f73d">More...</a><br/></td></tr>
<tr class="separator:a6264c3383db3384184cb13f5f282f73d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Exclusive prefix sum operations</div></td></tr>
<tr class="memitem:acd75d5aad2d1385bcbe15517011800e8"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#acd75d5aad2d1385bcbe15517011800e8">ExclusiveSum</a> (T input, T &amp;output)</td></tr>
<tr class="memdesc:acd75d5aad2d1385bcbe15517011800e8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element.  <a href="#acd75d5aad2d1385bcbe15517011800e8">More...</a><br/></td></tr>
<tr class="separator:acd75d5aad2d1385bcbe15517011800e8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a1fd276abbe08f7031a0333bf5c98c2f5"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a1fd276abbe08f7031a0333bf5c98c2f5">ExclusiveSum</a> (T input, T &amp;output, T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a1fd276abbe08f7031a0333bf5c98c2f5"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a1fd276abbe08f7031a0333bf5c98c2f5">More...</a><br/></td></tr>
<tr class="separator:a1fd276abbe08f7031a0333bf5c98c2f5"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a97e041cbdfdb4005ffc65b45c9276403"><td class="memTemplParams" colspan="2">template&lt;typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:a97e041cbdfdb4005ffc65b45c9276403"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a97e041cbdfdb4005ffc65b45c9276403">ExclusiveSum</a> (T input, T &amp;output, T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:a97e041cbdfdb4005ffc65b45c9276403"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a97e041cbdfdb4005ffc65b45c9276403">More...</a><br/></td></tr>
<tr class="separator:a97e041cbdfdb4005ffc65b45c9276403"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Exclusive prefix sum operations (multiple data per thread)</div></td></tr>
<tr class="memitem:ab8122c00c833f17c78af3d99dc76c5e8"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:ab8122c00c833f17c78af3d99dc76c5e8"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ab8122c00c833f17c78af3d99dc76c5e8">ExclusiveSum</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ab8122c00c833f17c78af3d99dc76c5e8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements.  <a href="#ab8122c00c833f17c78af3d99dc76c5e8">More...</a><br/></td></tr>
<tr class="separator:ab8122c00c833f17c78af3d99dc76c5e8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aa92950e9b763459fb7e63c8d047c94c2"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:aa92950e9b763459fb7e63c8d047c94c2"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#aa92950e9b763459fb7e63c8d047c94c2">ExclusiveSum</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], T &amp;block_aggregate)</td></tr>
<tr class="memdesc:aa92950e9b763459fb7e63c8d047c94c2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#aa92950e9b763459fb7e63c8d047c94c2">More...</a><br/></td></tr>
<tr class="separator:aa92950e9b763459fb7e63c8d047c94c2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ab6f9a0fa1ac832dcfb31bfacd7f1e22b"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:ab6f9a0fa1ac832dcfb31bfacd7f1e22b"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ab6f9a0fa1ac832dcfb31bfacd7f1e22b">ExclusiveSum</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:ab6f9a0fa1ac832dcfb31bfacd7f1e22b"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#ab6f9a0fa1ac832dcfb31bfacd7f1e22b">More...</a><br/></td></tr>
<tr class="separator:ab6f9a0fa1ac832dcfb31bfacd7f1e22b"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Exclusive prefix scan operations</div></td></tr>
<tr class="memitem:a2cd6dc7b523db630f4719b1b77df4db7"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:a2cd6dc7b523db630f4719b1b77df4db7"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a2cd6dc7b523db630f4719b1b77df4db7">ExclusiveScan</a> (T input, T &amp;output, T identity, ScanOp scan_op)</td></tr>
<tr class="memdesc:a2cd6dc7b523db630f4719b1b77df4db7"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element.  <a href="#a2cd6dc7b523db630f4719b1b77df4db7">More...</a><br/></td></tr>
<tr class="separator:a2cd6dc7b523db630f4719b1b77df4db7"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a1e09b0882138e34c23559c4d9a89a0d8"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:a1e09b0882138e34c23559c4d9a89a0d8"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a1e09b0882138e34c23559c4d9a89a0d8">ExclusiveScan</a> (T input, T &amp;output, const T &amp;identity, ScanOp scan_op, T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a1e09b0882138e34c23559c4d9a89a0d8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a1e09b0882138e34c23559c4d9a89a0d8">More...</a><br/></td></tr>
<tr class="separator:a1e09b0882138e34c23559c4d9a89a0d8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ab61af70303ff86b4aead54c1f83f0e30"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp , typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:ab61af70303ff86b4aead54c1f83f0e30"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ab61af70303ff86b4aead54c1f83f0e30">ExclusiveScan</a> (T input, T &amp;output, T identity, ScanOp scan_op, T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:ab61af70303ff86b4aead54c1f83f0e30"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#ab61af70303ff86b4aead54c1f83f0e30">More...</a><br/></td></tr>
<tr class="separator:ab61af70303ff86b4aead54c1f83f0e30"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Exclusive prefix scan operations (multiple data per thread)</div></td></tr>
<tr class="memitem:a5b21dfebcaf900cb516a06746dcb48b1"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </td></tr>
<tr class="memitem:a5b21dfebcaf900cb516a06746dcb48b1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a5b21dfebcaf900cb516a06746dcb48b1">ExclusiveScan</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], const T &amp;identity, ScanOp scan_op)</td></tr>
<tr class="memdesc:a5b21dfebcaf900cb516a06746dcb48b1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements.  <a href="#a5b21dfebcaf900cb516a06746dcb48b1">More...</a><br/></td></tr>
<tr class="separator:a5b21dfebcaf900cb516a06746dcb48b1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a1cda15ada4f880a7f428bd248c686710"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </td></tr>
<tr class="memitem:a1cda15ada4f880a7f428bd248c686710"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a1cda15ada4f880a7f428bd248c686710">ExclusiveScan</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], const T &amp;identity, ScanOp scan_op, T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a1cda15ada4f880a7f428bd248c686710"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a1cda15ada4f880a7f428bd248c686710">More...</a><br/></td></tr>
<tr class="separator:a1cda15ada4f880a7f428bd248c686710"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ab954851d7bd0fc7c3e2e16d31bfcb704"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:ab954851d7bd0fc7c3e2e16d31bfcb704"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ab954851d7bd0fc7c3e2e16d31bfcb704">ExclusiveScan</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], T identity, ScanOp scan_op, T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:ab954851d7bd0fc7c3e2e16d31bfcb704"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#ab954851d7bd0fc7c3e2e16d31bfcb704">More...</a><br/></td></tr>
<tr class="separator:ab954851d7bd0fc7c3e2e16d31bfcb704"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Inclusive prefix sum operations</div></td></tr>
<tr class="memitem:a570505ebdc51e2e47373fcb87e9a7d62"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a570505ebdc51e2e47373fcb87e9a7d62">InclusiveSum</a> (T input, T &amp;output)</td></tr>
<tr class="memdesc:a570505ebdc51e2e47373fcb87e9a7d62"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element.  <a href="#a570505ebdc51e2e47373fcb87e9a7d62">More...</a><br/></td></tr>
<tr class="separator:a570505ebdc51e2e47373fcb87e9a7d62"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a0de622762b27b095583770c66b905358"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a0de622762b27b095583770c66b905358">InclusiveSum</a> (T input, T &amp;output, T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a0de622762b27b095583770c66b905358"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a0de622762b27b095583770c66b905358">More...</a><br/></td></tr>
<tr class="separator:a0de622762b27b095583770c66b905358"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a39c481101dfa78b09374cab355712de9"><td class="memTemplParams" colspan="2">template&lt;typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:a39c481101dfa78b09374cab355712de9"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a39c481101dfa78b09374cab355712de9">InclusiveSum</a> (T input, T &amp;output, T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:a39c481101dfa78b09374cab355712de9"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a39c481101dfa78b09374cab355712de9">More...</a><br/></td></tr>
<tr class="separator:a39c481101dfa78b09374cab355712de9"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Inclusive prefix sum operations (multiple data per thread)</div></td></tr>
<tr class="memitem:a88ffea45e2a8b5e3abb6e4c4777e66ef"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:a88ffea45e2a8b5e3abb6e4c4777e66ef"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a88ffea45e2a8b5e3abb6e4c4777e66ef">InclusiveSum</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:a88ffea45e2a8b5e3abb6e4c4777e66ef"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements.  <a href="#a88ffea45e2a8b5e3abb6e4c4777e66ef">More...</a><br/></td></tr>
<tr class="separator:a88ffea45e2a8b5e3abb6e4c4777e66ef"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a553e70dd3e177545837438ded03b3bfd"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:a553e70dd3e177545837438ded03b3bfd"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a553e70dd3e177545837438ded03b3bfd">InclusiveSum</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a553e70dd3e177545837438ded03b3bfd"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a553e70dd3e177545837438ded03b3bfd">More...</a><br/></td></tr>
<tr class="separator:a553e70dd3e177545837438ded03b3bfd"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ae1a4a4dfbec4ec029dd6a8cce8b6c0a1"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:ae1a4a4dfbec4ec029dd6a8cce8b6c0a1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#ae1a4a4dfbec4ec029dd6a8cce8b6c0a1">InclusiveSum</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:ae1a4a4dfbec4ec029dd6a8cce8b6c0a1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#ae1a4a4dfbec4ec029dd6a8cce8b6c0a1">More...</a><br/></td></tr>
<tr class="separator:ae1a4a4dfbec4ec029dd6a8cce8b6c0a1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Inclusive prefix scan operations</div></td></tr>
<tr class="memitem:afb56064490291f37a712bcc3064ccbab"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:afb56064490291f37a712bcc3064ccbab"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#afb56064490291f37a712bcc3064ccbab">InclusiveScan</a> (T input, T &amp;output, ScanOp scan_op)</td></tr>
<tr class="memdesc:afb56064490291f37a712bcc3064ccbab"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element.  <a href="#afb56064490291f37a712bcc3064ccbab">More...</a><br/></td></tr>
<tr class="separator:afb56064490291f37a712bcc3064ccbab"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a60d6f1fb7780e25c4a1442128113270b"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp &gt; </td></tr>
<tr class="memitem:a60d6f1fb7780e25c4a1442128113270b"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a60d6f1fb7780e25c4a1442128113270b">InclusiveScan</a> (T input, T &amp;output, ScanOp scan_op, T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a60d6f1fb7780e25c4a1442128113270b"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a60d6f1fb7780e25c4a1442128113270b">More...</a><br/></td></tr>
<tr class="separator:a60d6f1fb7780e25c4a1442128113270b"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a7e1622b40fe73bdb6550f4dca2ae290a"><td class="memTemplParams" colspan="2">template&lt;typename ScanOp , typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:a7e1622b40fe73bdb6550f4dca2ae290a"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a7e1622b40fe73bdb6550f4dca2ae290a">InclusiveScan</a> (T input, T &amp;output, ScanOp scan_op, T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:a7e1622b40fe73bdb6550f4dca2ae290a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a7e1622b40fe73bdb6550f4dca2ae290a">More...</a><br/></td></tr>
<tr class="separator:a7e1622b40fe73bdb6550f4dca2ae290a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Inclusive prefix scan operations (multiple data per thread)</div></td></tr>
<tr class="memitem:afac59f8a498efb6a97f6c4a3b239576f"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </td></tr>
<tr class="memitem:afac59f8a498efb6a97f6c4a3b239576f"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#afac59f8a498efb6a97f6c4a3b239576f">InclusiveScan</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], ScanOp scan_op)</td></tr>
<tr class="memdesc:afac59f8a498efb6a97f6c4a3b239576f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements.  <a href="#afac59f8a498efb6a97f6c4a3b239576f">More...</a><br/></td></tr>
<tr class="separator:afac59f8a498efb6a97f6c4a3b239576f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a2885c51314333f4b98dacefcd7c918fc"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </td></tr>
<tr class="memitem:a2885c51314333f4b98dacefcd7c918fc"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#a2885c51314333f4b98dacefcd7c918fc">InclusiveScan</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], ScanOp scan_op, T &amp;block_aggregate)</td></tr>
<tr class="memdesc:a2885c51314333f4b98dacefcd7c918fc"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#a2885c51314333f4b98dacefcd7c918fc">More...</a><br/></td></tr>
<tr class="separator:a2885c51314333f4b98dacefcd7c918fc"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:abfce2a901cbcaac1852fef9f53d91a3a"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixOp &gt; </td></tr>
<tr class="memitem:abfce2a901cbcaac1852fef9f53d91a3a"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html#abfce2a901cbcaac1852fef9f53d91a3a">InclusiveScan</a> (T(&amp;input)[ITEMS_PER_THREAD], T(&amp;output)[ITEMS_PER_THREAD], ScanOp scan_op, T &amp;block_aggregate, BlockPrefixOp &amp;block_prefix_op)</td></tr>
<tr class="memdesc:abfce2a901cbcaac1852fef9f53d91a3a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs.  <a href="#abfce2a901cbcaac1852fef9f53d91a3a">More...</a><br/></td></tr>
<tr class="separator:abfce2a901cbcaac1852fef9f53d91a3a"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="ad0decc1ea510cd9c7df4cd380a26b1b1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::<a class="el" href="classcub_1_1_block_scan.html">BlockScan</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html" title="The operations exposed by BlockScan require a temporary memory allocation of this nested type for thr...">TempStorage</a> </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00262">262</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ab364204badf83769f9140d799cc188a7"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::<a class="el" href="classcub_1_1_block_scan.html">BlockScan</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00273">273</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a6264c3383db3384184cb13f5f282f73d"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::<a class="el" href="classcub_1_1_block_scan.html">BlockScan</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_scan_1_1_temp_storage.html" title="The operations exposed by BlockScan require a temporary memory allocation of this nested type for thr...">TempStorage</a> </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td><b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00284">284</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="a2cd6dc7b523db630f4719b1b77df4db7"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix max scan of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix max scan</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in those threads will be <code>INT_MIN, 0, 0, 2, ..., 124, 126</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00744">744</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a1e09b0882138e34c23559c4d9a89a0d8"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const T &amp;&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix max scan of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix max scan</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in those threads will be <code>INT_MIN, 0, 0, 2, ..., 124, 126</code>. Furthermore the value <code>126</code> will be stored in <code>block_aggregate</code> for all threads.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00795">795</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ab61af70303ff86b4aead54c1f83f0e30"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp , typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an exclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (block_aggregate &gt; old_prefix) ? block_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(INT_MIN);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide exclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveScan(</div>
<div class="line">            thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>INT_MIN, 0, 0, 2, ..., 124, 126</code>. The output for the second segment will be <code>126, 128, 128, 130, ..., 252, 254</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>126</code> in all threads after the first scan, assigned <code>254</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00888">888</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a5b21dfebcaf900cb516a06746dcb48b1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const T &amp;&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix max scan of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix max scan</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [INT_MIN,0,0,2], [2,4,4,6], ..., [506,508,508,510] }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00951">951</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a1cda15ada4f880a7f428bd248c686710"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">const T &amp;&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix max scan of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix max scan</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [INT_MIN,0,0,2], [2,4,4,6], ..., [506,508,508,510] }</code>. Furthermore the value <code>510</code> will be stored in <code>block_aggregate</code> for all threads.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01012">1012</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ab954851d7bd0fc7c3e2e16d31bfcb704"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>identity</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an exclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (block_aggregate &gt; old_prefix) ? block_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockLoad, BlockStore, and BlockScan for 128 threads, 4 ints per thread</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">cub::BlockLoad&lt;int*, 128, 4, BLOCK_LOAD_TRANSPOSE&gt;</a>   BlockLoad;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">cub::BlockStore&lt;int*, 128, 4, BLOCK_STORE_TRANSPOSE&gt;</a> BlockStore;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a>                             <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">union </span>{</div>
<div class="line">        <span class="keyword">typename</span> BlockLoad::TempStorage     load;</div>
<div class="line">        <span class="keyword">typename</span> BlockScan::TempStorage     scan;</div>
<div class="line">        <span class="keyword">typename</span> BlockStore::TempStorage    store;</div>
<div class="line">    } temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128 * 4)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">        BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide exclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage.scan).ExclusiveScan(</div>
<div class="line">            thread_data, thread_data, INT_MIN, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>INT_MIN, 0, 0, 2, 2, 4, ..., 508, 510</code>. The output for the second segment will be <code>510, 512, 512, 514, 514, 516, ..., 1020, 1022</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>510</code> in all threads after the first scan, assigned <code>1022</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">identity</td><td>Identity value </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01123">1123</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="acd75d5aad2d1385bcbe15517011800e8"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix sum of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, ..., 1</code>. The corresponding output <code>thread_data</code> in those threads will be <code>0, 1, ..., 127</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00335">335</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a1fd276abbe08f7031a0333bf5c98c2f5"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix sum of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data, block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, ..., 1</code>. The corresponding output <code>thread_data</code> in those threads will be <code>0, 1, ..., 127</code>. Furthermore the value <code>128</code> will be stored in <code>block_aggregate</code> for all threads. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00380">380</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a97e041cbdfdb4005ffc65b45c9276403"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an exclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total += block_aggregate;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveSum(</div>
<div class="line">            thread_data, thread_data, block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>1, 1, 1, 1, 1, 1, 1, 1, ...</code>. The corresponding output for the first segment will be <code>0, 1, ..., 127</code>. The output for the second segment will be <code>128, 129, ..., 255</code>. Furthermore, the value <code>128</code> will be stored in <code>block_aggregate</code> for all threads after each scan.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00465">465</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ab8122c00c833f17c78af3d99dc76c5e8"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix sum of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], ..., [508,509,510,511] }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00519">519</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aa92950e9b763459fb7e63c8d047c94c2"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an exclusive prefix sum of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).ExclusiveSum(thread_data, thread_data, block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], ..., [508,509,510,511] }</code>. Furthermore the value <code>512</code> will be stored in <code>block_aggregate</code> for all threads.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00574">574</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ab6f9a0fa1ac832dcfb31bfacd7f1e22b"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::ExclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an exclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total += block_aggregate;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockLoad, BlockStore, and BlockScan for 128 threads, 4 ints per thread</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">cub::BlockLoad&lt;int*, 128, 4, BLOCK_LOAD_TRANSPOSE&gt;</a>   BlockLoad;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">cub::BlockStore&lt;int*, 128, 4, BLOCK_STORE_TRANSPOSE&gt;</a> BlockStore;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a>                             <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">union </span>{</div>
<div class="line">        <span class="keyword">typename</span> BlockLoad::TempStorage     load;</div>
<div class="line">        <span class="keyword">typename</span> BlockScan::TempStorage     scan;</div>
<div class="line">        <span class="keyword">typename</span> BlockStore::TempStorage    store;</div>
<div class="line">    } temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128 * 4)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">        BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide exclusive prefix sum</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage.scan).ExclusiveSum(</div>
<div class="line">            thread_data, thread_data, block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>1, 1, 1, 1, 1, 1, 1, 1, ...</code>. The corresponding output for the first segment will be <code>0, 1, 2, 3, ..., 510, 511</code>. The output for the second segment will be <code>512, 513, 514, 515, ..., 1022, 1023</code>. Furthermore, the value <code>512</code> will be stored in <code>block_aggregate</code> for all threads after each scan.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00680">680</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="afb56064490291f37a712bcc3064ccbab"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix max scan of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix max scan</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in those threads will be <code>0, 0, 2, 2, ..., 126, 126</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01814">1814</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a60d6f1fb7780e25c4a1442128113270b"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix max scan of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix max scan</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>0, -1, 2, -3, ..., 126, -127</code>. The corresponding output <code>thread_data</code> in those threads will be <code>0, 0, 2, 2, ..., 126, 126</code>. Furthermore the value <code>126</code> will be stored in <code>block_aggregate</code> for all threads.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01864">1864</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a7e1622b40fe73bdb6550f4dca2ae290a"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename ScanOp , typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes one input element. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an inclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (block_aggregate &gt; old_prefix) ? block_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(INT_MIN);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide inclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveScan(</div>
<div class="line">            thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>0, 0, 2, 2, ..., 126, 126</code>. The output for the second segment will be <code>128, 128, 130, 130, ..., 254, 254</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>126</code> in all threads after the first scan, assigned <code>254</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01956">1956</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="afac59f8a498efb6a97f6c4a3b239576f"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix max scan of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix max scan</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,0,2,2], [4,4,6,6], ..., [508,508,510,510] }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l02016">2016</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a2885c51314333f4b98dacefcd7c918fc"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename ScanOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix max scan of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix max scan</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveScan(thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,0,2,2], [4,4,6,6], ..., [508,508,510,510] }</code>. Furthermore the value <code>510</code> will be stored in <code>block_aggregate</code> for all threads.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l02085">2085</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="abfce2a901cbcaac1852fef9f53d91a3a"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveScan </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">ScanOp&#160;</td>
          <td class="paramname"><em>scan_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using the specified binary <code>scan_op</code> functor. Each thread contributes an array of consecutive input elements. the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Supports non-commutative scan operators.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an inclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total = (block_aggregate &gt; old_prefix) ? block_aggregate : old_prefix;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockLoad, BlockStore, and BlockScan for 128 threads, 4 ints per thread</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">cub::BlockLoad&lt;int*, 128, 4, BLOCK_LOAD_TRANSPOSE&gt;</a>   BlockLoad;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">cub::BlockStore&lt;int*, 128, 4, BLOCK_STORE_TRANSPOSE&gt;</a> BlockStore;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a>                             <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">union </span>{</div>
<div class="line">        <span class="keyword">typename</span> BlockLoad::TempStorage     load;</div>
<div class="line">        <span class="keyword">typename</span> BlockScan::TempStorage     scan;</div>
<div class="line">        <span class="keyword">typename</span> BlockStore::TempStorage    store;</div>
<div class="line">    } temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128 * 4)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">        BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide inclusive prefix max scan</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage.scan).InclusiveScan(</div>
<div class="line">            thread_data, thread_data, <a class="code" href="structcub_1_1_max.html" title="Default max functor. ">cub::Max</a>(), block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>0, -1, 2, -3, 4, -5, ...</code>. The corresponding output for the first segment will be <code>0, 0, 2, 2, 4, 4, ..., 510, 510</code>. The output for the second segment will be <code>512, 512, 514, 514, 516, 516, ..., 1022, 1022</code>. Furthermore, <code>block_aggregate</code> will be assigned <code>510</code> in all threads after the first scan, assigned <code>1022</code> after the second scan, etc.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scan_op</td><td>Binary scan operator </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l02202">2202</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a570505ebdc51e2e47373fcb87e9a7d62"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix sum of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix sum</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, ..., 1</code>. The corresponding output <code>thread_data</code> in those threads will be <code>1, 2, ..., 128</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01381">1381</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a0de622762b27b095583770c66b905358"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix sum of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain input item for each thread</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data;</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix sum</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveSum(thread_data, thread_data, block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>1, 1, ..., 1</code>. The corresponding output <code>thread_data</code> in those threads will be <code>1, 2, ..., 128</code>. Furthermore the value <code>128</code> will be stored in <code>block_aggregate</code> for all threads. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01426">1426</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a39c481101dfa78b09374cab355712de9"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>output</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an inclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total += block_aggregate;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data = d_data[block_offset];</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide inclusive prefix sum</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveSum(</div>
<div class="line">            thread_data, thread_data, block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        d_data[block_offset] = thread_data;</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>1, 1, 1, 1, 1, 1, 1, 1, ...</code>. The corresponding output for the first segment will be <code>1, 2, ..., 128</code>. The output for the second segment will be <code>129, 130, ..., 256</code>. Furthermore, the value <code>128</code> will be stored in <code>block_aggregate</code> for all threads after each scan.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input item </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output item (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01512">1512</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a88ffea45e2a8b5e3abb6e4c4777e66ef"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix sum of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix sum</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveSum(thread_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [1,2,3,4], [5,6,7,8], ..., [509,510,511,512] }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01566">1566</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a553e70dd3e177545837438ded03b3bfd"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates an inclusive prefix sum of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockScan for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockScan::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute the block-wide inclusive prefix sum</span></div>
<div class="line">    <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">    <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage).InclusiveSum(thread_data, thread_data, block_aggregate);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [1,2,3,4], [5,6,7,8], ..., [509,510,511,512] }</code>. Furthermore the value <code>512</code> will be stored in <code>block_aggregate</code> for all threads.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">ScanOp</td><td><b>[inferred]</b> Binary scan operator type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01631">1631</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ae1a4a4dfbec4ec029dd6a8cce8b6c0a1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename BlockPrefixOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_scan.html">cub::BlockScan</a>&lt; T, BLOCK_THREADS, ALGORITHM &gt;::InclusiveSum </td>
          <td>(</td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>output</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T &amp;&#160;</td>
          <td class="paramname"><em>block_aggregate</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">BlockPrefixOp &amp;&#160;</td>
          <td class="paramname"><em>block_prefix_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor <code>block_prefix_op</code> is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs. Also provides every thread with the block-wide <code>block_aggregate</code> of all inputs. </p>
<p>The <code>block_prefix_op</code> functor must implement a member function <code>T operator()(T block_aggregate)</code>. The functor's input parameter <code>block_aggregate</code> is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from <em>lane</em><sub>0</sub> is applied as the block-wide prefix. Can be stateful.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates a single thread block that progressively computes an inclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// A stateful callback functor that maintains a running prefix to be applied</span></div>
<div class="line"><span class="comment">// during consecutive scan operations.</span></div>
<div class="line"><span class="keyword">struct </span>BlockPrefixOp</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Running prefix</span></div>
<div class="line">    <span class="keywordtype">int</span> running_total;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Constructor</span></div>
<div class="line">    __device__ BlockPrefixOp(<span class="keywordtype">int</span> running_total) : running_total(running_total) {}</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Callback operator to be entered by the first warp of threads in the block.</span></div>
<div class="line">    <span class="comment">// Thread-0 is responsible for returning a value for seeding the block-wide scan.</span></div>
<div class="line">    __device__ <span class="keywordtype">int</span> operator()(<span class="keywordtype">int</span> block_aggregate)</div>
<div class="line">    {</div>
<div class="line">        <span class="keywordtype">int</span> old_prefix = running_total;</div>
<div class="line">        running_total += block_aggregate;</div>
<div class="line">        <span class="keywordflow">return</span> old_prefix;</div>
<div class="line">    }</div>
<div class="line">};</div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, <span class="keywordtype">int</span> num_items, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockLoad, BlockStore, and BlockScan for 128 threads, 4 ints per thread</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">cub::BlockLoad&lt;int*, 128, 4, BLOCK_LOAD_TRANSPOSE&gt;</a>   BlockLoad;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">cub::BlockStore&lt;int*, 128, 4, BLOCK_STORE_TRANSPOSE&gt;</a> BlockStore;</div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan&lt;int, 128&gt;</a>                             <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan</span></div>
<div class="line">    __shared__ <span class="keyword">union </span>{</div>
<div class="line">        <span class="keyword">typename</span> BlockLoad::TempStorage     load;</div>
<div class="line">        <span class="keyword">typename</span> BlockScan::TempStorage     scan;</div>
<div class="line">        <span class="keyword">typename</span> BlockStore::TempStorage    store;</div>
<div class="line">    } temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Initialize running total</span></div>
<div class="line">    BlockPrefixOp prefix_op(0);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have the block iterate over segments of items</span></div>
<div class="line">    <span class="keywordflow">for</span> (<span class="keywordtype">int</span> block_offset = 0; block_offset &lt; num_items; block_offset += 128 * 4)</div>
<div class="line">    {</div>
<div class="line">        <span class="comment">// Load a segment of consecutive items that are blocked across threads</span></div>
<div class="line">        <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">        BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Collectively compute the block-wide inclusive prefix sum</span></div>
<div class="line">        <span class="keywordtype">int</span> block_aggregate;</div>
<div class="line">        <a class="code" href="classcub_1_1_block_scan.html#a982e1407d00b704c3046cd72c48acabb" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockScan</a>(temp_storage.scan).IncluisveSum(</div>
<div class="line">            thread_data, thread_data, block_aggregate, prefix_op);</div>
<div class="line">        __syncthreads();</div>
<div class="line"></div>
<div class="line">        <span class="comment">// Store scanned items to output segment</span></div>
<div class="line">        BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);</div>
<div class="line">        __syncthreads();</div>
<div class="line">    }</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the input <code>d_data</code> is <code>1, 1, 1, 1, 1, 1, 1, 1, ...</code>. The corresponding output for the first segment will be <code>1, 2, 3, 4, ..., 511, 512</code>. The output for the second segment will be <code>513, 514, 515, 516, ..., 1023, 1024</code>. Furthermore, the value <code>512</code> will be stored in <code>block_aggregate</code> for all threads after each scan.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">BlockPrefixOp</td><td><b>[inferred]</b> Call-back functor type having member <code>T operator()(T block_aggregate)</code> </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">output</td><td>Calling thread's output items (may be aliased to <code>input</code>) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">block_aggregate</td><td>block-wide aggregate reduction of input items (exclusive of the <code>block_prefix_op</code> value) </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">block_prefix_op</td><td><b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a block-wide prefix to be applied to all inputs. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l01744">1744</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<hr/>The documentation for this class was generated from the following file:<ul>
<li><a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a></li>
</ul>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:12 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
